[SOUND] Hi, everyone.
Welcome back to the heterogeneous parallel programming class.
We are at lecture 1.7, kernel-based parallel programming.
And we will be discussing multidimensional kernel configuration today.
The objective of this lecture is for you to
understand the use of multidimensional thread grids and in particular,
we're going to be discussing the use of multidimensional block
and thread indices, and we're also be discussing how we
can impact this map block and thread indices to
data indices to processes two dimensional pictures and so on.
And after this lecture you should be able to easily
generalize all the concepts into the processing of three dimensional
data structures using three dimensional blocks and grids.
This picture was shown earlier in the introductory lecture where we
assumed that the host has launched a thread of a kernel one.
And this generates a grid of threads.
And we assume that the user has specified that we're going to generate a
two- dimensional grid which is reflected
by the two-dimensional indices of the blocks.
And then we also assume that each block has three dimensional threads.
So each thread in this picture has three indices.
The convention that we're going to be using is that for a for a two dimensional
structure we're going to show x and y indexes, the first
index will be x, and the second index will be y.
And for a three dimensional structure, we're going to show
the indices in the order of x, y, and z.
So the first index will be x, second index will y, and the third index will be z.
So, where, here, we have a very simple example and now we're ready
to look at a more practical example. So we, we
see here that a here a two dimensional picture of
pixels. And we have 62 pixels in the y
dim, dimension, and 72 6 pixels in the x dimension.
So we will be calling this picture a 62 by 76 picture.
Which means that the first number in the, in the
configuration will be the number of pixels in the y dimension.
And the second number will be in the x dimension.
We assume that we're going to be using a 16 by 16 thread block.
16 threads in the y dimension, and 16
threads in the x dimension to process the picture.
Just like in the vector, addition example. We need to make sure that we have enough
threads in both y dimension and x dimension to process all the pixels.
So we're going to need to launch four thread blocks
in the y dimension to cover all the 62 pixels.
When we launch four thread blocks, we have 64 threads.
So we
have all the 62 pixels covered, and we also need to have 2 ex-, we need to have,
we need to, we will have two extra threads
in the y dimension that should not be doing anything.
And in the x dimension we're going to need to launch five thread blocks.
And that will give us 80 threads.
So we could, we would cover all the 76 pixels and we will have
four threads that should not do anything. And this will be reflected in the way we
write the kernel. Before we look at the kernel code,
I'd like to quickly review the the way C and C plus
plus lay out a two dimensional matrix, or two dimensional
array into the linear address space. In the modern
computers, all the memory data are stored in a linear address space
and even though the conceptually, we have a two dimensional array in C,
ultimately all the elements will be stored in a linear order and this is called the
layout and C adopts the tradition, convention of Row-Major Layout.
Which means that all the elements in a row are preserved in their relative position.
So all the neighboring elements in row zero
will be placed consecutively in the, memory, linearized memory.
And all the elements in row one will also be placed in consecutive locations.
So, in this particular case, whenever we need to access a,
a element using a linearized address. We can generate that,
address by multiplying the row index by the
width of the of the array, and plus the column index.
So, for example, you wanted to access element
M21. We have a row index of two, and a column
index of one.
So we can multiply the, the, the row index 2 by the width, which is 4.
And then we plus the column index, which is 1, and this gives us 9.
So if you look at the top of the picture,
we show the linearized index for all the M elements.
The first four are zero, one, two, three, and then
the next ones are four, five, six, seven. And so element
nine is actually the original M two, one. And
this will be used for accessing all the dynamically
allocated arrays in in the CUDA kernels.
And so the, even though, this is
familiar to all of you.
It's very important that you keep this in mind
when you look at any of the CUDA kernels.
And, for those of you who came from Fortran background.
Fortran adopts a, nat, column nature layout convention, where all the
elements in the, in the row, column, will have their relative positions preserved.
So, in this particular case,
if you look at the C, layout. The, the adjacent
elements in a column are actually laid out in a, strided way,
that are, have mult- so that they are with a throw width away from each other.
Whereas, if you had a Fortran program, all these elements
in the same column will actually be placed in adjacent locations.
Now that we, we understand how, CUDA, or C and C
Plus Plus linearize the a two dimensional array.
Now we're ready to look at a, piece of kernel code.
So here is a picture kernel, which is a very, very simple kernel that what
a, multiply the pixel, every pixel value in a picture
by 2.0.
So, this is essentially magnifying the value of
all the pic, pixel, pixels in a picture.
So, here we show that the the kernel accepts four arguments.
The first one is pointer to the input picture.
The second one is pointer to the output picture.
And the third one is the the number of elements in the in
the row, in the x direction. And the fourth one is the
number of elements in the in the vertical or
y dimension. So when we
determine which pixel a particular thread, should, process.
We have a, y index generation, and x index generation.
The y index generation generates a row index.
Remember that row index selects one of the rows in that, in the, y dimension.
So that's why we're using the y indices.
So we will just use our familiar
pattern of blockIdx.y time blockDim.y plus threadIdx.y.
And this ensures that every possible row will
have a, a thread that that covers that row.
And then we have the same pattern for the x dimension.
And this is pretty much identical to the vector addition case, where every element
in the explanation will have a thread to cover that particular pixel value.
So, now that we have generated appropriate row index and
column index.
And all the threads can now use the
row index and column index to do the processing.
We can first test whether a whether the thread is within the valid range.
So we will first test if row index is less than m.
So m is the number of pixels in the y dimension.
So the row index needs
to be within that range.
And also, we will test whether the column index is less than n.
N is the number of pixels in the x dimension.
So whenever both conditions are satisfied, we know that a thread is, we think, the
valid range of a two-dimensional area where there, there are valid pixels.
So now the thread can proceed to to multiply the
the input pixel by two and assign that to the output.
So, here, you show, you see the linearized addressing where, the thread
is going to use row multiplied by n, which is number of pixels in each row.
So this is the width, plus, column index.
And, you're, we're also going to use exactly the
same expression to write into the output matrix.
So, or output array.
So this will allow every thread to pick the correct its assigned row and
column index position pixel, multiply by 2, and write that into the output array.
So here we show the host code that that launches the pic, picture kernel.
So, we assume that the host host in the host code, the variable
n, host the number of pixels in the y dimension and
the n shows the host the number of pixels in the x dimension [COUGH].
In our example of
62 times 76, m will be equal to 62 and n will be equal to [COUGH]
76.
So now we were going to launch the kernel.
And this is how we make sure that all the kernels,
excuse me, will have enough threads.
The kernel will be executed by enough threads to
cover both the y dimension and the x dimension.
So, we take, the, n value and we are
doing a seeding function, essentially a seeding function of
n divided by 16, this n minus the familiar
expression of n minus 1 divided by 16 plus 1.
[COUGH] We
do the same thing for x. So we have m minus 1 divided by 16 plus 1.
[COUGH] Excuse me. [SOUND] And
then, with each block, we're going to just have 16 threads in the x dimension.
16 threads in the y, dimension. And one thread in the z dimension.
So this gives us a two dimensional grid, a two dimensional
block structure to cover the picture given.
And then we, we use the configuration to launch the kernel,
and then we call the PictureKernel with the DimGrid and DimBlock.
Now we give the the four arguments.
The pointer to the input picture, pointer to the output picture, m and n.
So this will ensure that the kernel will have enough
threads to cover both dimensions and then you'll, we will be able to
in the kernel we'll be able to test whether each thread should be
acted, acting on a pixel or the thread is outside the valid
dimension, so that, valid range, so that the thread should not take any action.
So this picture shows a little bit more
analysis of what actually happens at run time.
So we launched four thread blocks by five thread blocks.
So we actually end up launching 20 thread blocks.
And then we're actually going to have four different situations,
as far as the execution of the thread blocks are concerned.
The first case is in the upper left corner.
Where all the threads will have both their y
index, and x index, to be in the valid range.
So, all the threads within these three by four, 12, all these 12 thread
blocks will have all their threads, fully executing, and fully processing
the picture.
And the second one, that's the one on the upper right side.
These threads will find their y indices within the valid range, but
x indices in the not in the valid range, well some of
the x indices will be in the valid range, and some of
the x x in, indices will not be in the valid range.
So, in this case, some of the threads
in that in, in the, on the right end, in these thread blocks will not execute.
And then we have the third, situation where, all the threads in
the x dimension will find their x indices in the valid range.
But some of the threads in the y dimension
will find their index to be outside the valid range.
So in this,
in the third case, we will have this thread blocks will have some of their
threads in the bottom of their of the thread block not processing any pixels.
And finally, we have the fourth case.
Where the threads, some of the threads in the x, will find their x index
to be in the, outside the valid range. And some of the
threads will find their y index to be outside the valid range.
And in fact some of the threads in the very corner, the lower right corner, will
find their, both of their x and y indices outside the value range.
So these, you know, situations will all happen.
And this is an important picture for you to remember.
When we just begin to discuss the performance implications.
For now, it suffices for you to see, to remember that whenever we do,
we have a process, a two dimensional picture and we use boundary check.
To to check whether each thread has its y index or x
index within the valid range then we will end up with these four possible
situations, or four possible cases.
And this brings us to the conclusion of this lecture.
And for those of you who would like to understand more about the material,
I would like to encourage you to read sections 4.1 and 4.2 of the textbook.
Thank you.